#include <iostream>
#include <chrono>
#include <thread>
#include <cuda_runtime.h>
#include <cuda_profiler_api.h>

#include "helper_timer.h"
#include "helper_cuda.h"

__global__ void increment_kernel(int *g_data, int inc_value)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    g_data[idx] = g_data[idx] + inc_value;
}

bool correct_output(int *data, const int n, const int x)
{
    for (int i = 0; i < n; i++)
        if (data[i] != x)
        {
            printf("Error! data[%d] = %d, ref = %d\n", i, data[i], x);
            return false;
        }

    return true;
}

int main()
{
    // ==== profile time spent in cpu ====
    StopWatchInterface *timer = nullptr;
    sdkCreateTimer(&timer);
    sdkResetTimer(&timer);

    sdkStartTimer(&timer);
    std::this_thread::sleep_for(std::chrono::seconds(1));
    sdkStopTimer(&timer);

    printf("time spent by cpu: %.2f ms\n", sdkGetTimerValue(&timer));

    // ==== profile time spent in gpu ====
    int n = 16 * 1024 * 1024;
    int nbytes = n * sizeof(int);
    int value = 26;

    // allocate host memory
    int *a = 0;
    checkCudaErrors(cudaMallocHost((void **)&a, nbytes));
    memset(a, 0, nbytes);

    // allocate device memory
    int *d_a = 0;
    checkCudaErrors(cudaMalloc((void **)&d_a, nbytes));
    checkCudaErrors(cudaMemset(d_a, 255, nbytes));

    // set kernel launch configuration
    dim3 threads = dim3(512, 1);
    dim3 blocks = dim3(n / threads.x, 1);

    // create cuda event handles
    cudaEvent_t start, stop;
    checkCudaErrors(cudaEventCreate(&start));
    checkCudaErrors(cudaEventCreate(&stop));

    checkCudaErrors(cudaDeviceSynchronize());
    float gpu_time = 0.0f;

    // start profiling time spent in gpu
    checkCudaErrors(cudaProfilerStart());
    cudaEventRecord(start, 0);
    cudaMemcpyAsync(d_a, a, nbytes, cudaMemcpyHostToDevice, 0);
    increment_kernel<<<blocks, threads, 0, 0>>>(d_a, value);
    cudaMemcpyAsync(a, d_a, nbytes, cudaMemcpyDeviceToHost, 0);
    cudaEventRecord(stop, 0);
    checkCudaErrors(cudaProfilerStop());

    checkCudaErrors(cudaDeviceSynchronize()); // we can either wait for device to complete computing,
                                              // or do sth else in the meantime.

    checkCudaErrors(cudaEventElapsedTime(&gpu_time, start, stop));
    printf("time spent executing by the GPU: %.2f ms\n", gpu_time);

    if (correct_output(a, n, value))
    {
        printf("program executed successfully!\n");
    }
}